{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "# 代码生成：VM 基础"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 2,
   "metadata": {},
   "outputs": [],
   "source": [
    "import set_env"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "metadata": {},
   "outputs": [],
   "source": [
    "import tvm\n",
    "import tvm.testing\n",
    "from tvm import te\n",
    "from tvm.script import tir as T, ir as I\n",
    "\n",
    "import numpy as np"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 4,
   "metadata": {},
   "outputs": [],
   "source": [
    "def run_jit(fapi, check):\n",
    "    for target in [\"llvm\", \"stackvm\"]:\n",
    "        if not tvm.testing.device_enabled(target):\n",
    "            continue\n",
    "        f = tvm.driver.build(fapi, target=target)\n",
    "        s = f.get_source()\n",
    "        check(f)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "def test_stack_vm_basic():\n",
    "    a = tvm.nd.array(np.zeros(10, dtype=\"float32\"))\n",
    "\n",
    "    @tvm.register_func\n",
    "    def tvm_call_back_get_shape(shape0):\n",
    "        print(shape0)\n",
    "        assert shape0 == a.shape[0]\n",
    "\n",
    "    n = te.size_var(\"n\")\n",
    "    Ab = tvm.tir.decl_buffer((n,), \"float32\")\n",
    "    stmt = tvm.tir.Evaluate(tvm.tir.call_packed(\"tvm_call_back_get_shape\", Ab.shape[0]))\n",
    "\n",
    "    mod = tvm.IRModule.from_expr(\n",
    "        tvm.tir.PrimFunc([Ab], stmt).with_attr(\"global_symbol\", \"print_shape\")\n",
    "    )\n",
    "\n",
    "    run_jit(mod, lambda f: f(a))\n",
    "\n",
    "\n",
    "@tvm.register_func\n",
    "def tvm_stack_vm_print(*x):\n",
    "    print(x)\n",
    "\n",
    "\n",
    "def test_stack_vm_loop():\n",
    "    dtype = \"int64\"\n",
    "    n = te.size_var(\"n\")\n",
    "    Ab = tvm.tir.decl_buffer((n,), dtype)\n",
    "    i = te.size_var(\"i\")\n",
    "\n",
    "    ib = tvm.tir.ir_builder.create()\n",
    "    A = ib.buffer_ptr(Ab)\n",
    "    with ib.for_range(0, n - 1, \"i\") as i:\n",
    "        A[i + 1] = A[i] + 1\n",
    "        ib.emit(tvm.tir.call_packed(\"tvm_stack_vm_print\", i))\n",
    "\n",
    "    stmt = ib.get()\n",
    "    mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([Ab], stmt).with_attr(\"global_symbol\", \"ramp\"))\n",
    "    a = tvm.nd.array(np.zeros(10, dtype=dtype))\n",
    "\n",
    "    def check(f):\n",
    "        f(a)\n",
    "        np.testing.assert_equal(a.numpy(), np.arange(a.shape[0]))\n",
    "\n",
    "    run_jit(mod, check)\n",
    "\n",
    "\n",
    "def test_stack_vm_cond():\n",
    "    dtype = \"int64\"\n",
    "    n = te.size_var(\"n\")\n",
    "    Ab = tvm.tir.decl_buffer((n,), dtype)\n",
    "\n",
    "    ib = tvm.tir.ir_builder.create()\n",
    "    A = ib.buffer_ptr(Ab)\n",
    "    with ib.for_range(0, n - 1, \"i\") as i:\n",
    "        with ib.if_scope(tvm.tir.EQ(i, 4)):\n",
    "            A[i + 1] = A[i] + 1\n",
    "        with ib.else_scope():\n",
    "            A[i + 1] = A[i] + 2\n",
    "\n",
    "    stmt = ib.get()\n",
    "    mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([Ab], stmt).with_attr(\"global_symbol\", \"test\"))\n",
    "\n",
    "    def check(f):\n",
    "        a = tvm.nd.array(np.zeros(10, dtype=dtype))\n",
    "        f(a)\n",
    "        y = np.arange(a.shape[0]) * 2\n",
    "        y[5:] -= 1\n",
    "        np.testing.assert_equal(a.numpy(), y)\n",
    "\n",
    "    run_jit(mod, check)\n",
    "\n",
    "\n",
    "def test_vm_parallel():\n",
    "    dtype = \"int64\"\n",
    "    n = te.size_var(\"n\")\n",
    "    Ab = tvm.tir.decl_buffer((n,), dtype)\n",
    "    i = te.size_var(\"i\")\n",
    "    ib = tvm.tir.ir_builder.create()\n",
    "    A = ib.buffer_ptr(Ab)\n",
    "    with ib.for_range(0, n, \"i\", kind=\"parallel\") as i:\n",
    "        A[i] = A[i] + 1\n",
    "    stmt = ib.get()\n",
    "    mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([Ab], stmt).with_attr(\"global_symbol\", \"test\"))\n",
    "\n",
    "    def check(f):\n",
    "        a = tvm.nd.array(np.zeros(10, dtype=dtype))\n",
    "        f(a)\n",
    "        np.testing.assert_equal(a.numpy(), np.ones(a.shape[0]))\n",
    "\n",
    "    run_jit(mod, check)\n",
    "\n",
    "\n",
    "def test_codegen_decl_buffer():\n",
    "    \"\"\"The codegen should accept DeclBuffer nodes in its input\"\"\"\n",
    "\n",
    "    @I.ir_module\n",
    "    class mod:\n",
    "        @T.prim_func\n",
    "        def kernel(A_data: T.handle(\"float32\")):\n",
    "            T.func_attr({\"global_symbol\": \"kernel\"})\n",
    "            A_buf = T.decl_buffer([256], dtype=\"float32\", scope=\"global\", data=A_data)\n",
    "\n",
    "    target = tvm.target.Target(\"stackvm\")\n",
    "    stackvm_codegen = tvm.get_global_func(\"target.build.stackvm\")\n",
    "    stackvm_codegen(mod, target)\n"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "py312x",
   "language": "python",
   "name": "python3"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 3
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython3",
   "version": "3.12.3"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 2
}
